Ascend Send Ordering

导言

之前学习 Data-as-Flag 时，我们讨论过把 flag 塞进数据块：每个 512 B 块中，480 B 是有效数据，32 B 用来证明“这一块已经到达”。现在希望减少这部分控制开销，改成“前面 relaxed order 发送，最后一个 strong order”。

这个改动的核心，是把“每块都带证明”改成“用队列末尾的有序操作，为前面一批传输提供完成证明”。 但必须解释清楚：最后一个是什么、保证覆盖哪条队列、接收方如何得知完成，以及何时可以复用 buffer。本文沿实际代码逐层回答。

先把需求说准确

这里的 relaxoder 通常写作 Relaxed Order；所选 SHMEM 头文件使用 Relax Order（RO）。strongorder 写作 Strong Order（SO）。两者描述通信工作请求之间的执行约束。

“32 B 无效数据”更准确地说是32 B 非业务数据：它不参与模型计算，却承担同步证明。只有找到新的正确性保证，才可以拿掉它。

最重要的结论有三条：

1. RO 数据 + SO 收尾可以减少逐块 flag 开销。 RO 允许前面的独立写入保留较大的执行自由度；后面的 SO 建立一个必须等待前序完成的边界。
2. 最后一个必须按 PE/QP 分别考虑。 PE 是通信参与者，通常对应一个 rank；QP 是该参与者的一条通信队列上下文。QP0 的 SO 不替 QP1 等待。
3. SO 本身不等于接收方已经得知完成。 代码可以把 SO 放在一个接收方轮询的标志写入上，也可以把最后一条数据设成 SO，由发送方等待完成通知，再通知接收方。

本文固定分析以下源码，后续行号均对应这些版本：

来源	固定版本	用途
ascend_deepep_0909	f319d4c6fa41b9433f432e06732d00b9ad8b6bc4	用户给定分支的 kernels/elastic_combine.cpp
该版本锁定的 cann/shmem 子模块	38a47e53f37c8814c5b9f4067adfc0be0a3691ba	include/device/gm2gm/engine/shmem_device_udma.h 的 API 契约

固定版本的 combine 代码和配套 SHMEM 头文件共同构成实现依据。这里说明的是 Ascend950 UDMA 路径；没有在 NPU 上执行本次性能或正确性实验，也不把它泛化为所有昇腾型号、MTE 或 RoCE 后端的统一保证。

从装箱理解顺序

把发送方想成仓库 A，接收方想成仓库 B。业务数据是货物，flag 是一张“可以开始使用”的通知单。

• 逐块带 flag：每个封好的箱子里都有自己的到货证明。接收方验证某个箱子的证明，就可以用那个箱子的货物。
• RO + SO 标志：前面的货物可以灵活运输，但最后那张通知单必须等本队列前面的货都送完，才能执行发送。接收方看到通知单，才开始使用这一批货物。

这个比喻的关键是“最后那张通知单有强制的先后约束”。如果它只是普通消息，通知单完全可能先到，货物还在路上。


自绘概念插图：收尾者负责等待同一队列的前序工作；它不是所有队列的全局裁判。图中的货物只是通信写入的比喻，不表示物理网络按此形状布线。

四个时刻

先区分一次异步发送中的四件事：

时刻	发生了什么	能立即推出什么
提交	软件把工作请求交给引擎	引擎有工作要做；不能据此读远端数据
传输完成	对应请求已完成，数据到达目标 PE	要按 API 契约判断哪些源 buffer 可复用
接收方观察完成	接收侧通过信号等待协议看到本轮标志	在顺序、代际和可见性条件满足时，开始读相应数据
协议回收	队列空间、缓冲区和本轮信号可以复用	下一轮不会覆盖尚未使用的数据或证据

put_nbi 中 NBI 指非阻塞接口。函数返回通常只是提交完成，不能直接等同于传输完成。 配套头文件第 455–470 行明确要求使用 aclshmemx_udma_qp_quiet(pe, qp_idx) 等待，才能复用源数据或依赖远端可见性。QP Put 契约

三种 order

所选头文件第 22–49 行定义了三种 order。先把 NO 和 RO 分开，因为这会直接影响正确性。

名称	编码	在本文协议中的含义
No Order，NO	0b000	不受后续 SO 对前序 RO/SO 的这项保证覆盖
Relax Order，RO	0b101	可作为后续 SO 要等待的前序请求；不能仅凭提交次序认定彼此的完成次序
Strong Order，SO	0b110	在同一 PE/QP 上，先等前面的 RO/SO WQE 完成，再执行自己

一个可用的例子是：

同一 PE/QP 的提交次序：D0(RO) → D1(RO) → D2(RO) → F(SO)
允许的现象：D2 比 D0 先完成
必须满足：D0、D1、D2 都完成之后，F 才执行
接收侧：通过正确的等待接口观察 F，再读取 D0、D1、D2

D0、D1、D2 表示互不重叠的业务数据写入，F 表示单独的完成标志写入。这个规则只建立必要的边界，并不要求前面所有数据严格逐条串行传输。

默认 NO 不能直接套用

此版本 SHMEM 默认配置是 NO。不能省略配置、发一批默认 put_nbi，最后补 SO，就宣称前面的请求都被覆盖。需要显式指定前序 RO 或 SO，并核对目标 PE 和 QP。ordering 契约，第 36–49 行

不要把这里的 RO/SO 与 C++ std::memory_order_relaxed、release/acquire 或 PCIe Relaxed Ordering 直接画等号。它们作用的对象和保证范围不同；理解时可以借用“发布数据”的直觉，落实时必须使用当前通信后端的契约。

读代码前认识这些对象

UDMA 在这里是 UnifiedBus 相关的数据搬运引擎；Ascend C kernel 生成工作请求，由引擎执行远端写入。软件填的是“搬运任务说明”，payload 不等于这份任务说明。

对象	源码名称或教学符号	形状与类型	谁产生、谁使用	生命周期
发送特征行	x / local_hidden，记作 D_i	[H]，此路径支持 BF16；字节数 hidden_bytes = H × x_element_bytes	模型计算或本地 reduce 产生，UDMA 读取	对应传输完成前保持稳定
接收特征行	remote_hidden / symmetric_buffer 中的 hidden 区	[H]；行距为 hidden_row_stride	UDMA 写，接收侧 reduce 读	写入到消费结束
工作请求	WQE，Work Queue Entry	本实现一个 WRITE WQE 为 64 B	AIV/SIMT 构造，UDMA 取走执行	发布后到队列回收；64 B 不是每行新增的远端 payload
发送队列	SQ，Send Queue	WQE 环形数组；位置为 head % depth	kernel 生产，UDMA 消费	跨批次使用，受 credit 限制
队列上下文	HierarchyUdmaQueueState / QP	地址、目标信息、initial_head/head/depth 等	SHMEM 初始化，kernel 更新	按 peer 和 QP 分开管理
完成记录	CQE，Completion Queue Entry；存放于 CQ	控制记录，包含 owner、status、entry index	引擎产生，发送侧 poll 读取	用于等待、错误处理、队列回收
收尾信号	local_signal / remote_signal，记作 F	[1] uint64_t，逻辑写入 8 B	本地准备值 1，远端等待值 1	必须区分本轮与上一轮
提交暂存	udma_wqe_buffer	UB 中 128 × 64 = 8192 B	defer 填写，submit 发布	一个提交批次内
暂存描述符	valid_flags	UB 中 [token_count] uint8_t	Stage2 的 reduce 产生，SIMT 消费	当前 tile；不是每个网络数据块的 32 B flag
路由参数	destination_y/publisher_y/token_capacity	标量整数	tiling 和路由逻辑提供	算出发送、接收行位置，不是 payload

GM 是 Global Memory，全局内存地址空间；UB 是 Unified Buffer，核内暂存空间；AIV 指向量计算核心；SIMT 表示多线程按同一程序执行的编程方式。本文只跟踪通信与控制对象，不展开专家模型内部矩阵计算。

队列还有一个 doorbell（门铃）：软件写寄存器告知引擎“新任务已经准备好了”。它用于触发取任务，不是接收完成标志。

配置对象本身是四个字段：

// 依次是：是否生成 CQE、保留事件位、fence 位、order 编码。
aclshmemx_udma_op_config_t { cqe, se, fence, odr };

cqe 决定是否产生完成记录，odr 决定顺序约束。RO 可以生成 CQE，SO 也可以不生成 CQE。 fence 是另一个独立字段；本次三个配置都把它设为 0，顺序语义来自 odr，不能看见 0 就说“没有同步”。

方案一：RO 数据后追加 SO 标志

这种机制把完成证明从每个数据块挪到一个通信批次末尾。业务数据保持原样写入，额外发一个小标志；接收方等待标志后再消费这一批数据。

概念层是批次完成通知；通用机制是同一队列上 RO payload → SO signal；具体实现层包括本文件的直接 P2P 分支，以及 hierarchical Stage1。两者共享远端信号证明，但 CQE 和回收策略不同。

和之前的Data-as-Flag 文章比较时，要限定到其中 32 B flag + 480 B payload 的 TFF 布局，不能把它当成所有 Data-as-Flag 方案的固定成本。新方案可能省下逐块 pack/unpack 和控制字节，但增加批次边界等待、信号管理和队列管理；接收粒度也可能从块级变为批次级。


自绘五视图，依据 combine f319d4c6 的 P2P 路径和 Stage1。A 看逐块开销如何转为批次信号；B 看收益与等待粒度的交换；C 看初始化、提交、等待和回收；D 看源数据保活与接收信号；E 看业务数据与 8 B 信号走不同对象。图中两种 CQE 回收分支分别标注，不能混用。

P2P 如何写

第 200–208 行的三个配置就是一个很好的入门入口。下面保留原字段值和次序：

inline constexpr aclshmemx_udma_op_config_t
    kElasticCombineUdmaDataNoCqeConfig{
        0U, 0U, 0U, ACLSHMEMX_UDMA_ODR_RO};
inline constexpr aclshmemx_udma_op_config_t
    kElasticCombineUdmaDataCqeConfig{
        1U, 0U, 0U, ACLSHMEMX_UDMA_ODR_RO};
inline constexpr aclshmemx_udma_op_config_t
    kElasticCombineUdmaCompletionCqeConfig{
        1U, 0U, 0U, ACLSHMEMX_UDMA_ODR_SO};

它们分别用于：

1. 批内数据：RO，不生成 CQE，使用 defer_action 暂存请求。
2. 提交批次的最后一条数据：仍是 RO，生成 CQE，使用 submit_action 真正提交。源码按最多 128 条 WQE 分批，或者在数据耗尽时提交。
3. 该 peer/QP 的最终标志：SO，生成 CQE，写入一个 uint64_t 的 1。

因此，“一次 submit 的最后一条”和“最终 SO 标志”是两个不同边界。不要看到 DataCqeConfig 就把它当成 SO。数据提交第 2745–2857 行

第 2871–2881 行的最后一次 Put 是实际 SO 收尾：

aclshmemx_udma_qp_put_nbi<
    uint64_t, PIPE_MTE3,
    kElasticCombineUdmaCompletionCqeConfig>(
    completion_signal,
    completion_signal,
    reinterpret_cast<__ubuf__ uint64_t*>(udma_wqe_buffer),
    1U,
    source_rank,
    qp_slot,
    EVENT_ID5);

逐项读它：

• uint64_t 是发送元素类型，1U 表示 1 个元素，即 8 B，不是 1 B。
• 第一个 completion_signal 是对称内存目的地址，由 source_rank 指定目标 PE；第二个是当前 rank 的本地源地址。两个指针写成一样，不等于拷贝回自身：SHMEM 会根据目标 PE 处理远端对称地址。
• source_rank 在 combine 中是原 token 所属的目标 rank；变量名字来自 token 的来处，不代表它是当前发送进程。
• qp_slot 必须与前面这批数据使用的 QP 相同。此直接 P2P 路径当前只用 QP0。
• udma_wqe_buffer 暂存工作请求；PIPE_MTE3 表示发布 WQE 使用的管线，不能因此把远端 payload 搬运统称为普通 MTE DataCopy。
• EVENT_ID5 是相关本地硬件同步事件编号，不是远端业务 flag。

发送前，第 2589–2623 行把本地“其他发布者的接收槽”清 0，把自身发布者对应的源槽设 1，然后进行核间和 rank 间同步。即使某个 peer 没有 payload，仍发送末尾标志，否则接收方固定等待所有 peer 时可能永远等不到。标志发送与接收，第 2863–2920 行

P2P 的完整流程

下面规范化了 P2P 的通信路径，保留所有相关提交和等待条件；它不是可直接编译的 Ascend C kernel。requests[p] 是路由解析后要发给 peer p 的请求列表，每项有本地源、对称目的地址和字节数；hidden 与可选的 FP32 top-k weight 都展开为独立请求。W 是 rank 数，r 是当前 rank，scratch 是 8192 B UB WQE 暂存。路由数学不参与 order 证明，故作为输入而不冒充省略的通信步骤。

# 前置：SHMEM、对称窗口、W 个 peer 的 QP0 已初始化；每条队列有单一提交者；映射区可供目标 UDMA 访问。
# 所有源数据已由计算管线发布到 UDMA 可读取的内存，并保持稳定。
for publisher in 0 .. W-1:
    local_signal[publisher, 0] = (1 if publisher == r else 0)
wait_local_signal_stores_complete()
sync_local_cores()
sync_all_ranks()                         # 所有接收槽先完成初始化

for p in 0 .. W-1:
    if p == r:
        copy_self_requests_and_wait()
        continue
    state = new_empty_submit_state()
    for i in 0 .. len(requests[p])-1:
        req = requests[p][i]
        final_in_batch = (state.pending_count == 127)
        final_in_peer = (i == len(requests[p])-1)
        if final_in_batch or final_in_peer:
            put_bytes(req, p, qp=0, RO, CQE=1, submit(state), scratch)
        else:
            put_bytes(req, p, qp=0, RO, CQE=0, defer(state), scratch)
    assert state.pending_count == 0
    put_uint64(remote_slot=p.signal[r,0], source=local_signal[r,0],
               count=1, peer=p, qp=0, SO, CQE=1)  # 空数据 peer 也执行

for publisher in 0 .. W-1:
    if publisher != r:
        wait_until(local_signal[publisher,0] == 1)
sync_local_cores()                       # 分摊等待的核汇合
for p in 0 .. W-1:
    if p != r:
        qp_quiet(p, 0)                  # 排空自己的出站请求，完成资源管理
sync_local_cores()
consume_ready_partial_rows_in_reduce_epilogue()
# 下一轮还须遵守调度、窗口与信号生命周期契约；不得并发重置仍在使用的槽。

接收端等待的是远端内存里已经写好的值，发送端 quiet 等待的是自己的出站队列完成。此实现先由所有本地等待者观察入站标志，再执行出站 quiet；这不是要求把 quiet 放到每次 payload Put 之后。

Stage1 为什么没有 CQE

hierarchical Stage1 同样追加 8 B SO 信号，但使用 HierarchyMarkUdmaStrongOrderNoCqe。第 379–391 行先建立普通 WRITE，再将它改为 SO 且关闭 CQE；业务请求由 HierarchyWriteUdmaWqe 默认设为 RO。

对应的闭环是：

对一个目标 peer 的 QP0：
    检查 SQ 剩余容量，整个未回收区间不能越过允许界限
    并行生成本 step 的 RO 数据 WQE
    在 step 边界追加 8 B SO/no-CQE 信号 WQE
    发布 WQE，再写 doorbell
接收端：
    等待该 step 所有来源的 signal == 1
    读取对应特征行并执行 reduce
所有参与者结束本轮后：
    完成必要内存同步、核间同步、rank 间 barrier
    将 Stage1 队列 tail 更新为 head，回收已完成请求

这里 no-CQE 不等于不等完成。HierarchyPrepareStage1Queue 用容量检查防止未完成 WQE 被覆盖；退出时的 rank barrier 配合此前的信号消费，为 HierarchyReclaimStage1Queue 回收队列提供完成依据。不能把 tail = head 单独复制到一个普通异步发送循环。Stage1 容量检查与回收，第 1484–1504、1697–1712、2520–2534 行

这一变体减少 CQE 生成与轮询，把回收成本移到更粗的协议同步点；代价是严格的队列容量约束和更强的全轮生命周期耦合。移植工作归通信 kernel 与 SHMEM 集成层所有，不能只修改模型侧张量布局。

方案二：最后一条数据改成 SO

hierarchical Stage2 更接近“最后一个数据发 strong order”的字面意思：先构造一批默认 RO 的业务 WQE，再把每条非空 QP 的最后一个改成 SO，并打开 CQE。

HierarchyStage2SubmitVf 第 779–781 行按有效路由序号分流：

const uint32_t route_order = atomicAdd(
    &batch_state->valid_route_count, 1U);
const uint32_t qp_slot =
    (route_order & 3U) == 3U ? 1U : 0U;

这是每四个有效路由约三个给 QP0、一个给 QP1 的分配方式。atomicAdd 返回旧值并递增计数器，为并行线程分配路由序号；线程调度决定哪些具体 token 获得这些序号，因此不要把它理解成固定的“第 4 个 token 一定去 QP1”。

随后第 794–805 行执行：

asc_threadfence();
__syncthreads();
if (thread_id < kElasticCombineStage2UdmaQpCount) {
    __ubuf__ HierarchyUdmaQueueState* queue =
        &batch_state->queues[thread_id];
    if (queue->head != queue->initial_head) {
        HierarchyMarkUdmaStrongOrderCqe(queue, queue->head - 1U);
    }
}

这里 head 是下一处可分配的位置，initial_head 是本批开始位置。所以：

• head != initial_head：这条 QP 本批至少有一个 WQE。
• head - 1：这条 QP 本批最后一个已分配的 WQE。
• asc_threadfence() 和 __syncthreads()：先完成本地线程之间的描述符发布与汇合，再修改末条属性；它们不是远端通信完成等待。

真正的保证分别建立在两条队列中：

QP0：A(RO) → B(RO) → C(SO + CQE)
QP1：D(RO) → E(SO + CQE)

C 等待 QP0 中更早的请求后才执行；E 对 QP1 做同样的事。接收方不能靠检查 C 的某个数据值是否变化来替代完成等待，因为 C 是业务数据，不是协议定义的完成标志。Stage2 WQE 构造与末条修改，第 745–805 行


自绘五视图，依据 combine f319d4c6 的 Stage2。A 看每条队列末尾各有一个 SO；B 看减少逐请求通知与批次等待的交换；C 看先敲两条队列门铃、再等待两条 CQ；D 看 quiet 之后才发布接收通知；E 看 UB 有效位、GM 数据和完成控制流。Stage2 最后仍有独立信号，不是整个协议“零 flag”。

从描述符到引擎

HierarchyWriteUdmaWqe 在第 332–335 行把 RO 编码写进描述符。HierarchyMarkUdmaStrongOrderCqe 在第 357–366 行清除旧 order 位，再写 SO 与 CQE 位。用教学变量表示为：

words = SQ[(head mod depth)]             # 实际由字节地址与 wqe_size 定位
word0 = words[0]
word0 = word0 AND NOT(0b111 << 16)        # 清旧 order，不误伤其他字段
word0 = word0 OR ((0b110 OR (1 << 5)) << 16)
words[0] = word0                         # SO 在 16..18 位，CQE 位在 21 位

words 是一个 64 B WQE 的八个 uint64_t 视图；word0 是其中首个 64 位字，AND/OR/NOT 是位运算。head/depth/wqe_size 来自队列上下文。这是固定头文件布局的解释，不是建议初学者手工照抄这些数字。 源文件第 230–249 行用 static_assert 对照 SHMEM 的 opcode、结构大小、order 和 CQE 常量；升级 SHMEM 必须重新验证布局和语义。

描述符写好后，还没有自动交给硬件。HierarchyFlushAndRingUdmaQueue 第 561–600 行做两件关键工作：

1. 对新写的 WQE 区间执行 dcci_cachelines，处理环形队列回绕时的两段区间，使通信引擎能够读取正确的描述符。
2. 使用 st_dev(queue->head, doorbell_addr, 0) 写门铃，然后更新队列 bookkeeping；若本批有 CQE，还递增完成计数。

这条路径借助 SIMT 直接构造 GM 中的 WQE，属于对 SHMEM 内部数据结构的深度集成。公开 API 路径和手工 SQ 路径的完成计数管理不能混着用；本文件第 440–442 行也专门说明了为何自行轮询 CQ。原始 WQE 与提交路径，第 312–600 行

Stage2 的完整流程

这里 T 是每个目标的 token 容量，t 是 token 编号，y_dst 与 y_pub 分别映射 destination_y 与 publisher_y，S 是 hidden_row_stride，H_bytes 是有效行字节数。local_base/remote_base 是已建立映射的对称窗口基址；send_offset/recv_offset 分别是 Stage1 归约结果区和 Stage2 接收区偏移。所有地址表达式都在访问物理 GM 数据，并不创建新的张量副本。

# 单个远端 peer、单个 tile；输入是已写好的归约结果与 valid_flags。
# QP0/QP1 均已初始化，计算到通信的本地事件已完成。
检查两条 SQ 的 credit；必要时轮询已有 CQE 以回收容量
从 SHMEM 上下文读取每条 QP 的 initial_head 和 head
valid_route_count = 0

并行遍历 tile 内的 token t:
    if valid_flags[t - token_begin] == 0:
        continue
    send_row = y_dst * T + t
    recv_row = y_pub * T + t
    src = local_base + send_offset + send_row * S
    dst = remote_base + recv_offset + recv_row * S
    n = atomic_fetch_add(valid_route_count, 1)
    q = 1 if n mod 4 == 3 else 0
    h = atomic_fetch_add(queue[q].head, 1)
    在 queue[q] 的位置 h 写入 WRITE(src, dst, H_bytes, RO, CQE=0)

本地线程发布描述符并汇合
for q in [0, 1]:
    if queue[q].head != queue[q].initial_head:
        将该队列 head-1 对应 WQE 改成 SO、CQE=1
本地线程再次发布描述符并汇合

for q in [0, 1]:                         # 先提交两条，保留重叠机会
    if 该队列本批非空:
        刷新新 WQE 区间的缓存
        写该队列 doorbell
        更新 head 与预计 CQE 计数
for q in [0, 1]:
    if 该队列本批非空:
        轮询到目标 CQE 计数
        检查 owner、status、substatus；超时或异常按失败处理
        更新 CQ tail、门铃和 SQ tail

# 同一个 step 有多个 tile 时，先完成该 step 的所有 tile。
用 HierarchyPublishSignal 向目标发布 step 信号 1
接收方用 HierarchyWaitForSignal 等待这个 step 信号
接收方消费对应 GM 特征行并归约
在整轮退出 barrier 后，才进入允许复用本轮信号的下一轮

HierarchyFlushRemoteQueues 第 1507–1559 行明确先 ring 两条 QP，再逐条 quiet。Stage2 的 step 结束后，第 2395 行附近还调用 HierarchyPublishSignal：它经 aclshmem_ptr 取得远端地址，使用 DataCopyPad 写 8 B 信号，并等待本地 MTE3 完成。

因此，这个方案省掉的是每条数据中的 flag，以及部分细粒度完成通知；接收通知仍然存在。 因果链变成：

各 QP 最后一条 SO 数据完成
    → 发送侧 quiet 收到成功完成证明
    → 发布 step 信号
    → 接收方等待 step 信号
    → 消费本 step 数据

它适合能够承担 tile/step 完成等待、需要批量构造 WQE 的路径。新成本包括 CQE 轮询、credit 管理、直接操作内部结构的维护负担，以及等待后再通知的固定延迟。没有同机同形状实验，不能声称它一定比 SO 标志法更快。

能省多少，又不能省什么

不要把“信号值大小”“工作区槽位大小”和“网络实际流量”混为一谈。

设 P 为有效 payload 字节数；旧 TFF 布局每个 512 B 块容纳 480 B payload。这里只比较业务数据和逻辑标志，不含链路头、重传、硬件最小事务粒度：

旧方案的块数 B = ceil(P / 480)
旧方案数据区字节 = 512 × B
单队列、单批 SO 标志法的逻辑传输字节 = P + 8
多批多队列：逻辑标志字节 = 8 × 实际发送的标志次数

例如 P = 480 × 1000 = 480000 B：旧布局需要 512000 B，其中 flag 是 32000 B；若新方案只用一个 SO 标志，逻辑上是 480008 B。这只能说明这个假设下的编码开销变化，不是实际吞吐提升比例。32/512 是线速容量占比 6.25%，32/480 是相对有效 payload 的额外开销约 6.67%，两个分母不要混用。

代码还体现出三个独立成本：

• 8 B 逻辑信号：sizeof(uint64_t)，只表示 Put 请求搬的元素字节数。
• 512 B 槽距：kElasticCombineUdmaCompletionSignalBytes = 512；P2P 信号地址按 publisher_rank × 2 + qp_slot 再乘这个槽距计算。即使目前只用 QP0，物理寻址仍保留两 QP 布局。不能把 8 B 写入宣称为只占 8 B 工作区，也不能反向断言每次一定发 512 B 网络 payload。
• 64 B WQE 和 8192 B UB 暂存：它们是请求描述和批处理工作区，不是每个远端数据块携带的 32 B flag。

实际线上的事务对齐、协议头和最小粒度，需要目标硬件的计数器或规格进一步核实。Stage1 的信号区域也有自己的按 step/source 编排，不能套用 P2P 的 512 B 槽距。信号地址第 602–617 行

收益要连同同步一起测

路径	完成证明	主要可减少项	新增或保留成本
前文 TFF 布局	每块内嵌 flag 与数据的原子关系	独立通知路径；支持更细的消费粒度	逐块控制字节、pack/unpack、平台原子粒度依赖
P2P：RO 数据 + SO 标志	接收端观察同一 QP 的 SO 标志	逐块 flag 和相关打包	末尾标志、按批等待、出站 quiet、信号槽
Stage1：RO 数据 + SO/no-CQE 标志	接收信号与全轮 barrier	逐块 flag、CQE 生成和轮询	队列容量上限、退出同步与回收耦合
Stage2：末条 SO 数据 + CQE	每条 QP quiet 后发布 step 信号	逐块 flag、逐请求 CQE	CQ 轮询、step 信号、内部 WQE 布局适配

如需评估内存峰值，应在同一时刻累计：模型状态、活跃输入输出、保存的激活、算子暂存、通信 buffer、分配器余量和运行时余量。本优化直接影响的主要是通信布局、请求暂存和控制状态；不会自动减少模型参数或 hidden 数量。可以比较旧通信区中的 512 × ceil(P/480) 与新通信区的 payload、所有信号槽、队列和暂存，不能用上述逻辑流量公式直接代替整机峰值显存公式。

建议在固定硬件、CANN/SHMEM 版本、dtype、hidden、token 分布、rank 数、QP 数和批次大小的条件下，比较端到端 combine 时间、有效 payload 吞吐、P50/P99、WQE 构造时间、ring 时间、quiet 时间、接收等待时间、CQE 数以及通信区峰值。源码已有 kProfileUdmaBytes、kProfileSimtWqeBuildCycles、kProfileUdmaQuietCycles 等字段；其中 bytes 是软件计数，不能冒充链路实际字节计数。

改造时按这个顺序检查

1. 先锁定后端和版本。 README 指定的该 combine 路径是 Ascend950 节点内 UDMA、BF16 hidden、FP32 top-k weight；参与 NPU 需满足相同 scale-up domain 等运行条件。配套实现要求显式 QP API 使用 direct mode，ACLSHMEM_RELAY_SUPPORT=OFF；qp_idx 必须小于初始化的 QP 数，单个请求上限为 256 MB。更换到 A2/A3、MTE、其他 RDMA 或其他 SHMEM 版本，应重新建立契约，不能照抄 RO/SO 位编码。运行条件、QP API 平台分支
2. 画出每条 PE/QP 的请求串。 标出 RO 数据、SO 收尾、CQE 点和接收通知。确认前序不是默认 NO，并且同一批 payload 之间没有需要额外排序的重叠写入。
3. 把生产和发布接起来。 计算结果写入 GM 后，UDMA 才能读取；WQE 全部写好并完成必要缓存操作后，才能敲门铃。本地 SetFlag/WaitFlag、SIMT 同步、WQE cache 操作与远端 RO/SO 各自解决不同问题。
4. 明确最后一个的身份。 是独立 signal，还是最后一条数据？若是数据，接收方通过什么机制得知它已完成？每条非空 QP 是否都有收尾？
5. 处理空批次和槽位复用。 P2P 的空 peer 仍发 signal；Stage2 的空 QP 不应等待不存在的 CQE，但 step 级通知仍须满足接收协议。旧的 1 必须在正确的代际边界清零，或重新设计显式 generation 协议。原代码的常量 1 不等于通用多轮并发安全。
6. 保留消费结束和回收依据。 本地传输 quiet 可以结束源读取，但不能单独证明远端业务已消费完目标 buffer。远端复用还须等待消费者完成。Stage1 的 tail = head 依赖退出 barrier，不可提前。
7. 用能打破错误假设的实验验收。 除数值对照外，覆盖单条、多条、0 token、两 QP 不均衡、批次 127/128/129 边界、队列回绕、多轮复用、故意延迟其中一条队列。出现超时或 CQE 错误不能视作完成成功。

初学实现时可以先用配套 SHMEM 的显式 QP API验证协议，必要时每条 immediate 请求都保留 CQE；性能定位确实指向提交开销后，再考虑 defer/submit 或 SIMT 直接构造 WQE。这样每一步都能验证“谁为谁等、谁通知谁”，而不是一次同时改布局、队列和回收策略。

一句话复述

“前面 RO、最后 SO”是把同步证明从每块数据移到同一通信队列的收尾点：SO 等前面的 RO/SO 完成，接收方还必须通过明确的信号或通知协议才能安全消费。

在这份代码中，P2P 和 Stage1 用 8 B SO 标志收尾；Stage2 把每条 QP 的最后一条数据改成 SO + CQE，等待后再发 step 信号。它减少逐块 32 B flag 的潜在开销，但没有删除完成证明、队列管理和缓冲区生命周期。

参考资料

• 用户指定的分支源码：入口链接，分支会继续变化。
• 固定 combine f319d4c6：常量、WQE、P2P、Stage1/Stage2、接收与回收路径。
• 固定 SHMEM API 38a47e53：order 范围、NO 默认值、defer/submit、QP Put 和 quiet 契约。
• 固定项目 README：构建与运行平台边界。
• Data-as-Flag Communication：TFF、DC、SP 及 NVIDIA 对照的原文与论文证据。本文仅承接其中的 32 B TFF 布局，不重新宣称所有 Data-as-Flag 都有这项成本。

证据与实验边界

本文的 RO/SO 收尾方案由固定代码和配套 API 契约支撑，未确认其具有独立的方法论文，因此没有给它配造一对论文实验图。插图是教学示意；字节例子是注明假设的计算；本次只验证文档、源码对应关系和图形，不声称完成 NPU 编译、运行或性能验证。